﻿namespace RayDen.RayEngine.IntersectionEngines.OpenCL {
    internal partial class Kernels {


        public static string LuxBVHKernel2L = @"
typedef struct {
	float x, y, z;
} Point;

typedef struct {
    uint bvhOffset;
    uint index;
} Primitive;

typedef struct {
	float x, y, z;
} Vector;

typedef struct {
	Point o;
	Vector d;
	float mint, maxt;
} Ray;

typedef struct {
	float Distance;
	float U, V; // Barycentric coordinates of the hit point
	unsigned int Index;
} RayHit;

typedef struct {
	uint v[3];
    float d[3];
    uint count;
} TLHit;

typedef struct {
	unsigned int v[3];
} Triangle;

typedef struct {
	Point pMin, pMax;
} BBox;

typedef struct {
	BBox bbox;
	unsigned int primitive;
	unsigned int skipIndex;
} GpuBVHNode;

void TriangleIntersect(
		const float4 rayOrig,
		const float4 rayDir,
		const float minT,
		float *maxT,
		unsigned int *hitIndex,
		float *hitB1,
		float *hitB2,
		const unsigned int currentIndex,
		__global Point *verts,
		__global Triangle *tris,
        const unsigned int triangleOffset ) {

	// Load triangle vertices
	__global Point *p0 = &verts[tris[triangleOffset + currentIndex].v[0]];
	__global Point *p1 = &verts[tris[triangleOffset + currentIndex].v[1]];
	__global Point *p2 = &verts[tris[triangleOffset + currentIndex].v[2]];

	float4 v0 = (float4) (p0->x, p0->y, p0->z, 0.f);
	float4 v1 = (float4) (p1->x, p1->y, p1->z, 0.f);
	float4 v2 = (float4) (p2->x, p2->y, p2->z, 0.f);

	// Calculate intersection
	float4 e1 = v1 - v0;
	float4 e2 = v2 - v0;
	float4 s1 = cross(rayDir, e2);

	const float divisor = dot(s1, e1);
	if (divisor == 0.f)
		return;

	const float invDivisor = 1.f / divisor;

	// Compute first barycentric coordinate
	const float4 d = rayOrig - v0;
	const float b1 = dot(d, s1) * invDivisor;
	if (b1 < 0.f)
		return;

	// Compute second barycentric coordinate
	const float4 s2 = cross(d, e1);
	const float b2 = dot(rayDir, s2) * invDivisor;
	if (b2 < 0.f)
		return;

	const float b0 = 1.f - b1 - b2;
	if (b0 < 0.f)
		return;

	// Compute _t_ to intersection point
	const float t = dot(e2, s2) * invDivisor;
	if (t < minT || t > *maxT)
		return;

	*maxT = t;
	*hitB1 = b1;
	*hitB2 = b2;
	*hitIndex = currentIndex;
}

int BBoxIntersectPO(
		const float4 rayOrig, const float4 invRayDir,
		const float mint, const float maxt,
		const float4 pMin, const float4 pMax,float *t0, float *t1) {
	const float4 l1 = (pMin - rayOrig) * invRayDir;
	const float4 l2 = (pMax - rayOrig) * invRayDir;
	const float4 tNear = fmin(l1, l2);
	const float4 tFar = fmax(l1, l2);

	*t0 = max(max(max(tNear.x, tNear.y), max(tNear.x, tNear.z)), mint);
    *t1 = min(min(min(tFar.x, tFar.y), min(tFar.x, tFar.z)), maxt);

	return (t1 > t0);
}

int BBoxIntersectP(
		const float4 rayOrig, const float4 invRayDir,
		const float mint, const float maxt,
		const float4 pMin, const float4 pMax) {
	const float4 l1 = (pMin - rayOrig) * invRayDir;
	const float4 l2 = (pMax - rayOrig) * invRayDir;
	const float4 tNear = fmin(l1, l2);
	const float4 tFar = fmax(l1, l2);

	float t0 = max(max(max(tNear.x, tNear.y), max(tNear.x, tNear.z)), mint);
    float t1 = min(min(min(tFar.x, tFar.y), min(tFar.x, tFar.z)), maxt);

	return (t1 > t0);
}
__kernel void Intersect(
		__global Ray *rays,
		__global RayHit *rayHits,
		__global Point *verts,
        __global Primitive *prims,
		__global Triangle *tris,
		const unsigned int triangleCount,
		const unsigned int nodeCount,
		__global GpuBVHNode *bvhTree,		
		const unsigned int rayCount,
        __global TLHit *tlHits) {
	// Select the ray to check
	const int gid = get_global_id(0);
	if (gid >= rayCount)
		return;

    if (tlHits[gid].count == 0)    
        return;

    for (int ii=1;ii<tlHits[gid].count;ii++) {
        int primIndex = tlHits[gid].v[ii];
        if (primIndex > 10000)
            continue;
        float dd = 0.f;//tlHits[gid].d[ii];
        if (primIndex == 0xffffffffu)
            continue;
        unsigned int nodeOffset = prims[primIndex].bvhOffset;
	    unsigned int triangleOffset = prims[primIndex].index;

	    float4 rayOrig,rayDir;
	    float minT, maxT;
	    {
		    __global float4 *basePtr =(__global float4 *)&rays[gid];
		    float4 data0 = (*basePtr++);
		    float4 data1 = (*basePtr);

		    rayOrig = (float4)(data0.x + dd*data0.w, data0.y + dd*data1.x, data0.z+dd*data1.y, 0.f);
		    rayDir = (float4)(data0.w, data1.x, data1.y, 0.f);

		    minT = data1.z;
		    maxT = data1.w;
	    }

	    float4 invRayDir = (float4) 1.f / rayDir;

	    unsigned int hitIndex = 0xffffffffu;
	    unsigned int currentNode = nodeOffset; // Root Node
	    float b1, b2;
	    unsigned int stopNode = bvhTree[nodeOffset].skipIndex; // Non-existent

	    float4 pMin, pMax, data0, data1;
	    __global float4 *basePtr;
	        while (currentNode < stopNode) {

		        basePtr =(__global float4 *)&bvhTree[currentNode];
		        data0 = (*basePtr++);
		        data1 = (*basePtr);

		        pMin = (float4)(data0.x, data0.y, data0.z, 0.f);
		        pMax = (float4)(data0.w, data1.x, data1.y, 0.f);

		        if (BBoxIntersectP(rayOrig, invRayDir, minT, maxT, pMin, pMax)) {
			        //const unsigned int triIndex = bvhTree[currentNode].primitive;
			        const unsigned int triIndex = as_uint(data1.z);

			        if (triIndex != 0xffffffffu) {
				        TriangleIntersect(rayOrig, rayDir, minT, &maxT, &hitIndex, &b1, &b2, triIndex, verts, tris ,triangleOffset);
	                   // rayHits[gid].Distance   = maxT;
	                   // rayHits[gid].U = b1;
	                   // rayHits[gid].V = b2;
	                   // rayHits[gid].Index = triangleOffset+hitIndex;
                        //break;
                    } 
                    else {
                        //hitIndex = 0xffffffffu;
                        //maxT = 10000.0;

                    }


			        currentNode++;
		        } else {
			        //bvhTree[currentNode].skipIndex;
			        currentNode = as_uint(data1.w);
		        }                
	        }
            if (rayHits[gid].Distance > maxT) {
                    rayHits[gid].Distance   = maxT;
	                rayHits[gid].U = b1;
	                rayHits[gid].V = b2;
	                rayHits[gid].Index = triangleOffset+hitIndex;
                }
    }

} 

__kernel void TLIntersect(
		__global Ray *rays,
		__global TLHit *rayHits,
		const unsigned int nodeCount,
		__global GpuBVHNode *bvhTree,
		const unsigned int rayCount) {
	// Select the ray to check
	const int gid = get_global_id(0);
	if (gid >= rayCount)
		return;

	float4 rayOrig,rayDir;
	float minT, maxT;
	{
		__global float4 *basePtr =(__global float4 *)&rays[gid];
		float4 data0 = (*basePtr++);
		float4 data1 = (*basePtr);

		rayOrig = (float4)(data0.x, data0.y, data0.z, 0.f);
		rayDir = (float4)(data0.w, data1.x, data1.y, 0.f);

		minT = data1.z;
		maxT = data1.w;
	}

	//float4 rayOrig = (float4) (rays[gid].o.x, rays[gid].o.y, rays[gid].o.z, 0.f);
	//float4 rayDir = (float4) (rays[gid].d.x, rays[gid].d.y, rays[gid].d.z, 0.f);
	//float minT = rays[gid].mint;
	//float maxT = rays[gid].maxt;

	float4 invRayDir = (float4) 1.f / rayDir;

	unsigned int hitIndex = 0xffffffffu;
	unsigned int currentNode = 0; // Root Node
	float b1, b2,t0, t1;
    unsigned int v[3], vc =0;
	unsigned int stopNode = bvhTree[0].skipIndex; // Non-existent

	float4 pMin, pMax, data0, data1;
	__global float4 *basePtr;
	while (currentNode < stopNode) {

		basePtr =(__global float4 *)&bvhTree[currentNode];
		data0 = (*basePtr++);
		data1 = (*basePtr);

		pMin = (float4)(data0.x, data0.y, data0.z, 0.f);
		pMax = (float4)(data0.w, data1.x, data1.y, 0.f);

		if (BBoxIntersectPO(rayOrig, invRayDir, minT, maxT, pMin, pMax, &t0, &t1)) {
			//const unsigned int triIndex = bvhTree[currentNode].primitive;
			const unsigned int triIndex = as_uint(data1.z);

			if (triIndex != 0xffffffffu && vc < 3) {
                vc+=1;
                rayHits[gid].v[vc] = triIndex;
                rayHits[gid].d[vc] = t0;
                rayHits[gid].count = vc;
            }
			currentNode++;
		} else {
			currentNode = as_uint(data1.w);
		}
	}

	//rayHits[gid].v[0] = v[0];
	//rayHits[gid].v[1] = v[1];
	//rayHits[gid].v[2] = v[2];
	//rayHits[gid].count = vc;

} 
";




        
    }
}
